/** File name: pipeline.cu
*** Author: Tran Van Hoang
*** Date: 27/5/2015
*** Description: This version is developed on the previous version (cellRiver2.cu).
***              However, the pileline technique is used here.
***		 Dividing cell networks into many small parts.
***/

#include <sys/unistd.h>
#include <pthread.h>
#include <stdlib.h>
#include <stdio.h>
#include <cellNetwork.h>
#include <view.h>
#include "data.cu"
#include "network.cu"
#include <Xlib.h>

#define ROUNDS 20
#define THREADS 3
#define T1	0
#define T2	1
#define T3	2
#define BLSIZE 512
#define PART_SIZE 512

///Displaying Window
#define WINDOW_TITLE "CellRiver Pipeline version"
int plotSize, NSize;
Colormap colormap;
Display *dsp;	
int screenNumber;
unsigned long white;
unsigned long black;	
Window win;
GC gc;	
long eventMask;
XEvent evt;

//Main declaration
void *threadStep(void * args);

typedef struct {

	int threadId;
	NodeState *dataIn, *dataOut;
	canaux *channelIn, *channelOut;
	int part;//storing index of the part being performed

	}data_thread;

data_thread  data[THREADS];

extern __global__ void compute(NodeState *tabi,NodeState *tabout, canaux *channels, int part);

extern __global__ void SetZero(NodeState *tabi, int node_number);

void displayState(NodeState *nodeStates)
{
	int density;

    	for (int j = 0; j < NODE_NUMBER; j++)
    	{
	    density = (int) (nodeStates[j].density * 100.0);
	    if (density > 0.0)
	    {
    	    	showCellDensity(Cells[j], dsp, colormap, win, gc, plotSize, density);
	    }
    	}
}

void dataMove()
{
     	NodeState *tmpState;
     	canaux *tmpChannel;

        // Exchanging data between threads
	tmpState = data[0].dataOut;
	data[0].dataOut = data[1].dataIn;
	data[1].dataIn = tmpState;

	tmpState = data[1].dataOut;
	data[1].dataOut = data[2].dataIn;
	data[2].dataIn = tmpState;

	tmpChannel = data[0].channelOut;
	data[0].channelOut = data[1].channelIn;
	data[1].channelIn = tmpChannel;
}

void init_data_thread(data_thread *dt, int id)
{
	int i,err,nblocs;
        err= 0;
        nblocs = NODE_NUMBER / BLSIZE + 1;
	dt->threadId = id;
	dt->part = 0;

	switch (id)
		{
		case T1 : 
			dt->dataIn = (NodeState*) malloc(sizeof (NodeState) * NODE_NUMBER);
			//dt->channels = (canaux*) malloc(sizeof (canaux) * NODE_NUMBER);
			dt->channelIn = channels_h;//channels_h comes from network.cu
			err = cudaMalloc(&dt->dataOut,  sizeof (NodeState) * NODE_NUMBER);
			err = cudaMalloc(&dt->channelOut,  sizeof (canaux) * NODE_NUMBER);
			initState(dt->dataIn, NODE_NUMBER);
			break;
		case T2 : 
			err = cudaMalloc(&dt->dataIn,  sizeof (NodeState) * NODE_NUMBER);
			err |= cudaMalloc(&dt->dataOut,  sizeof (NodeState) * NODE_NUMBER);
			err = cudaMalloc(&dt->channelIn,  sizeof (canaux) * NODE_NUMBER);
  			SetZero<<<nblocs, BLSIZE>>>(dt->dataIn, NODE_NUMBER);
  			SetZero<<<nblocs, BLSIZE>>>(dt->dataOut, NODE_NUMBER);
			
			break;
		case T3 : 
			cudaMalloc(&dt->dataIn, sizeof (NodeState) * NODE_NUMBER);
  			SetZero<<<nblocs, BLSIZE>>>(dt->dataIn, NODE_NUMBER);
			dt->dataOut = (NodeState*) malloc(sizeof (NodeState) * NODE_NUMBER);			
			for (i=0; i < NODE_NUMBER; i++) 
			{
			    //dt->dataOut[i].density = 0.0 ;
			}
			break;
		}
}

void free_data(data_thread dt)
{
        int id;
	id = dt.threadId; 

	switch (id)
		{
		case T1 : 
			free(dt.dataIn);
			cudaFree(dt.dataOut);
			//free(dt.channelIn);
			cudaFree(dt.channelOut);
			break;
		case T2 :
			cudaFree(dt.dataIn);
			cudaFree(dt.dataOut);
			cudaFree(dt.channelIn);
			break;
		case T3 : 
			cudaFree(dt.dataIn);
			free(dt.dataOut);
			break;
		}
}

void initWindow()
{
	plotSize = 1;
	NSize = 256;
	dsp = XOpenDisplay(NULL);	
	if (!dsp){ return;}
    	colormap = DefaultColormap(dsp, 0);

	screenNumber = DefaultScreen(dsp);
	white = WhitePixel(dsp, screenNumber);
	black = BlackPixel(dsp, screenNumber);	

	win = XCreateSimpleWindow(dsp,
			DefaultRootWindow(dsp),
			50, 50,// origin
			NSize*plotSize*3, NSize*plotSize*3,// size
			0, black,//border
			white//background
	);

	XStoreName(dsp, win, WINDOW_TITLE);
	XMapWindow(dsp, win);
	
	gc = XCreateGC(dsp, win, 0, NULL);	
	
	eventMask = StructureNotifyMask;
	XSelectInput(dsp, win, eventMask);
	do
	{
		XNextEvent(dsp, &evt);
	}while(evt.type != MapNotify);	
}

void event()
{
	eventMask = ButtonPressMask | ButtonReleaseMask;

	XSelectInput(dsp, win, eventMask);
	do{
		XNextEvent(dsp, &evt);
	}while(evt.type != ButtonRelease);
}

int main(int argc, char * argv[])
{
    	int maxsteps;
    	void *status;

    	if (argc > 1)
	   maxsteps = atoi(argv[1]);
    	else
           maxsteps = ROUNDS;

    	pthread_t *threads;
    	pthread_attr_t attr;

    	threads=(pthread_t*) malloc( sizeof(pthread_t) * THREADS);
    	pthread_attr_init(&attr);
    	pthread_attr_setdetachstate(&attr, PTHREAD_CREATE_JOINABLE);

    	for (int i=0; i < THREADS;i++)
    	{
    	    init_data_thread(&data[i], i);
    	}
    	//Initilizating window
    	initWindow();
    	int density;

	for (int j = 0; j < NODE_NUMBER; j++)
	{
	     density = (int) (data[0].dataIn[j].density * 100.0);
	     showCellDensity(Cells[j], dsp, colormap, win, gc, plotSize, density);
	}

	printf("\n-- Click on the window to continue!\n");

	event();

	int nbPart;///Number of parts
		
	nbPart = (NODE_NUMBER / PART_SIZE);

	if ((nbPart * PART_SIZE) < NODE_NUMBER)
	{
	    nbPart++;
	}

    	//Simulation
    	for (int step = 0; step < maxsteps; step++)
    	{
	    for (int p = 0; p < (nbPart + 2); p++)
	    {
            	for (int i = 0; i < THREADS; i++)
       	    	{
		    ///Determining what part the thread need to execute
		    data[i].part = p - i;
           	    pthread_create(&threads[i], NULL, threadStep, (void*) &(data[i]));
            	}

            	for (int i = 0; i < THREADS; i++)
            	{
           	    pthread_join(threads[i], &status);
           	    if(status);
                    //Error;
            	}
	     	///Exchanging data between threads
             	dataMove();
	    }

	    //Preparing for next step
	    // Swap out2 -> in0	    	
	    NodeState *tmpState;
	    tmpState = data[2].dataOut;
	    data[2].dataOut = data[0].dataIn;
	    data[0].dataIn = tmpState;

            displayState(data[0].dataIn);
    	}
	
    	for (int i = 0; i < THREADS; i++) 
    	{
	    free_data (data[i]);
    	}

    	cudaDeviceReset();
	
	printf("\n-- Click on window to exit!\n");

	event();

    	printf("-- Pipe threads %d had finished for size Node_number=%d, nbParts=%d Rounds=%d\n",THREADS, NODE_NUMBER, nbPart, maxsteps);

    	return 0;
}

__global__ void SetZero(NodeState *nodeStates, int node_number)
{
    	int i = blockDim.x * blockIdx.x + threadIdx.x;

    	if (i < node_number)
    	{
	   nodeStates[i].density = (float) 0.0;
    	}
}

void toGPU(NodeState *tin, NodeState *tout, canaux *channels_h, canaux *channels_d, int part)
{
	int below, above;
	
	below = PART_SIZE * part;
	above = below + PART_SIZE;

	if (above > NODE_NUMBER)
	{
	    above = NODE_NUMBER;
	}
	
	/// Copying the main part
	cudaMemcpy(tout + below, tin + below, (above - below) * sizeof(NodeState), cudaMemcpyHostToDevice);
	/// Copying neighbourhoods of each cell. Omitting copied ones.
	int nbNeighbour, nodeIn;
	for (int i = below; i < above; i++)
	{
	    nbNeighbour = channels_h[i].nbIn;	   
	    for (int j = 0; j < nbNeighbour; j++)
	    {
	    	nodeIn = channels_h[i].read[j].node;
		if (nodeIn < below || nodeIn > above)
		{
		    cudaMemcpy(tout + nodeIn, tin + nodeIn, sizeof(NodeState), cudaMemcpyHostToDevice);
		}
	    }
	}

	/// Copying the links
	cudaMemcpy(channels_d, channels_h, NODE_NUMBER * sizeof(canaux), cudaMemcpyHostToDevice);
}

void processOnGPU(NodeState *tin, NodeState *tout, canaux *channels, int part)
{	
	int nblocs;

	nblocs = PART_SIZE / BLSIZE;
	if ((nblocs * BLSIZE) < PART_SIZE)
	{
	    nblocs++;
	}

	//compute<<<nblocs, BLSIZE>>>(tin, tout, channels, part);
}

void fromGPU(NodeState *tin, NodeState *tout, int part)
{
	int part_size;

	part_size = PART_SIZE;		/// The case of the last part
	if ((PART_SIZE * (part + 1)) > NODE_NUMBER) 
	{
	    part_size = NODE_NUMBER - (PART_SIZE * part);
	}
	cudaMemcpy(tout + (part * PART_SIZE), tin + (part * PART_SIZE), part_size * sizeof(NodeState), cudaMemcpyDeviceToHost);	
        //displayState(tout);
}

void *threadStep(void *args)
{
    	int id;
    	data_thread * sdata;
    	sdata = (data_thread*) args;
    	id  =  sdata->threadId;
	int part = sdata->part;
	int nbPart;
		
	nbPart = (NODE_NUMBER / PART_SIZE);

	if ((nbPart * PART_SIZE) < NODE_NUMBER)
	{
	    nbPart++;
	}

	if (part >= 0 && part < nbPart)
	{
    	    switch (id) {
	      	case T1: toGPU(sdata->dataIn, sdata->dataOut, sdata->channelIn, sdata->channelOut, part); break;
	      	case T2: processOnGPU(sdata->dataIn, sdata->dataOut, sdata->channelIn, part); break;
	      	case T3: fromGPU(sdata->dataIn, sdata->dataOut, part); break;
	    }
	}

    	return 0;
}

__global__ void compute(NodeState *inStates, NodeState *outStates, canaux *channels, int part)
{
    	int idx = blockDim.x * blockIdx.x + threadIdx.x;

        /// Changing to real id of cell in network
	idx = (PART_SIZE * part) + idx; 

    	if (idx < NODE_NUMBER) // idx < PART_SIZE * (part+1)
    	{

	    int nbIn = channels[idx].nbIn;
	    int nodeIn;
	    float receive = (float) 0.0;

	    for (int i = 0; i < nbIn; i++)
	    {
	    	nodeIn = channels[idx].read[i].node;
	    	receive = receive + ((inStates[nodeIn].density / 2.0) / (float) channels[nodeIn].nbIn);
	    }

	    outStates[idx].density = (inStates[idx].density / 2.0) + receive;
    	}
}

